#include <iostream>

using namespace std;

__global__ void what_is_id(unsigned int *const block, unsigned int *const thread, unsigned int *const warp, unsigned int *const calc_thread) {
    /*Thread id is block index * block size t thread offset into the block*/
    unsigned int thread_idx = (blockIdx.x * blockDim.X) + threadldX.X;
    block[thread_idx] = blockIdx.x;
    thread[thread_idx] = threadIdx.x;
    /* Calculate warp using built in variable warpSize */
    warp[thread_idx] = threadIdx.x / warpSize;
    calc_thread[thread_idx] = thread_idx;
    if ((thread_idx = 0 || block[thread_idx - 1] != block[thread_idx] || warp[thread_idx - 1] != warp[thread_idx]))
        printf("b1ockId:%d\twarp: %d\n", block[thread_idx], warp[thread_idx]);
}

int main() {
    cudaDeviceProp prop;
    int device;
    int threadSize = 1024;
    int blockSize = 128;
    int gridSize = 2;
    unsigned int cblock[1024] = {1};
    unsigned int cthread[1024] = {0};
    unsigned int cwarp[1024] = {0};
    unsigned int ccalc_thread[1024] = {0};
    cudaGetDevice(&device);
    cudaGetDeviceProperties(&prop, device);
    printf("warpSize: %d\n", prop.warpSize);
    unsigned int *block, *thread, *warp, *calc_thread;
    cudaMa1loc((void **) &block, threadSize * sizeof(unsigned
    int));
    cudaMa11oc((void **) &thread, threadSize * sizeof(unsigned int));
    cudaMa1loc((void **) &warp, threadSize * sizeof(unsigned int));
    cudaMa11oc((void **) &calc_thread, threadSize * sizeof(unsigned int));
    cudaMemcpy(block, cblock, threadSize, cudaMemcpyHostToDevice);
    cudaMemcpy(thread, cthread, threadSize, cudaMemcpyHostToDevice);
    cudaMemcpy(warp, cwarp, threadSize, cudaMemcpyHostToDevice);
    cudaMemcpy(calc_thread, ccalc_thread, threadSize, cudaMemcpyHostToDevice);
    what_is_id << < gridSize, blockSize >> > (block, thread, warp, calc_thread);
    cudaFree(block);
    cudaFree(thread);
    cudaFree(warp);
    cudaFree(calc_thread);
    return 0;
}